struct ScreenBlendVertex
{
	float x, y, z;
	float rhw;
	uint color;
	float tu, tv;			// The texture coordinates
	float tuBack, tvBack;	//The back texture coordinates
};

__global__ void GetCurveFlag(unsigned char* surface, size_t pitch, uint width, uint height, uint* occu, uint* ocuValid, ABSSURF_SHARE_INFO asi);

__global__ void NewtonPolish(unsigned char*surfaceIn, unsigned char*surfaceOut, size_t pitch,
							uint pixelBits,
							uint* curveSource, float* SilResult, uint actualPoint, float* FuncData, int* PatchIndex,
							ABSSURF_SHARE_INFO asi, int degSize,
							int* DebugInfo);

__global__ void compactVertex(uint pixelBits, uint* curveSource, ScreenBlendVertex* CurvePoint, uint actualPoint, RENDER_PARAM* multiSilRender, bool isForeshorten, uint radius);

__global__ void GetCurveFlag(unsigned char* surface, size_t pitch, uint width, uint height, uint pixelBits, uint* occu, uint* ocuValid, ABSSURF_SHARE_INFO asi)
{	
	uint x = blockIdx.x*blockDim.x + threadIdx.x;
	uint y = blockIdx.y*blockDim.y + threadIdx.y;
	if(x>=width || y>=height)			return;
	uint pixelI = y*width+x;
    float* pixel = (float*)(surface + y*pitch) + 4*x;
	
	if(pixel[3]==0.0)			return;

	occu[pixelI] = ((x<<pixelBits)|y);
	ocuValid[pixelI] = 1;
}
__global__ void compactVertex(uint pixelBits, uint* curveSource, ScreenBlendVertex* CurvePoint, uint actualPoint, bool isForeshorten, uint radius)
{
	uint pointN = threadIdx.x + blockIdx.x*blockDim.x;
	if(pointN>=actualPoint)							return;

	uint pixelInfo = curveSource[pointN];
	uint pixelY = (pixelInfo<<(sizeof(uint)*8-pixelBits)>>(sizeof(uint)*8-pixelBits));
	uint pixelX = (pixelInfo>>pixelBits);

	//float xOffset, yOffset;
	//if(!isForeshorten)
	//{	
	//}
	//else
	//{

	//}

	uint vertColor = 0xffff0000;
	uint pInd = 6*pointN;

	CurvePoint[pInd].x = pixelX - radius;
	CurvePoint[pInd].y = pixelY - radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 0.0f;
	CurvePoint[pInd].tv = 0.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;
	pInd++;

	CurvePoint[pInd].x = pixelX - radius;
	CurvePoint[pInd].y = pixelY + radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 0.0f;
	CurvePoint[pInd].tv = 1.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;
	pInd++;

	CurvePoint[pInd].x = pixelX + radius;
	CurvePoint[pInd].y = pixelY - radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 1.0f;
	CurvePoint[pInd].tv = 0.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;
	pInd++;
	
	CurvePoint[pInd].x = pixelX + radius;
	CurvePoint[pInd].y = pixelY - radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 1.0f;
	CurvePoint[pInd].tv = 0.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;
	pInd++;

	CurvePoint[pInd].x = pixelX - radius;
	CurvePoint[pInd].y = pixelY + radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 0.0f;
	CurvePoint[pInd].tv = 1.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;
	pInd++;
	
	CurvePoint[pInd].x = pixelX + radius;
	CurvePoint[pInd].y = pixelY + radius;
	CurvePoint[pInd].z = 0.01;
	CurvePoint[pInd].rhw = 1.0f;
	CurvePoint[pInd].color = vertColor;
	CurvePoint[pInd].tu = 1.0f;
	CurvePoint[pInd].tv = 1.0f;
	CurvePoint[pInd].tuBack = CurvePoint[pInd].x;
	CurvePoint[pInd].tvBack = CurvePoint[pInd].y;

}

__global__ void NewtonPolish(unsigned char*surfaceIn, unsigned char*surfaceOut, size_t pitch,
							uint pixelBits,
							uint* curveSource, float* SilResult, uint actualPoint, float* FuncData, int* PatchIndex,
							ABSSURF_SHARE_INFO asi, int degSize,
							int* DebugInfo)
{
	//uint pointN = threadIdx.x + blockIdx.x*blockDim.x;
	//if(pointN>=actualPoint)							return;

	//uint pixelInfo = curveSource[pointN];
	//uint pixelY = (pixelInfo<<(sizeof(uint)*8-pixelBits)>>(sizeof(uint)*8-pixelBits));
	//uint pixelX = (pixelInfo>>pixelBits);
	//float* pixelIn = ((float*)(surfaceIn+pixelY*pitch) + pixelX*4);

	//float screenX = 2.0f*(pixelX+0.5f)/cRenderParam.renderWidth-1.0f;
	//float screenY = 1.0f-2.0f*(pixelY+0.5f)/cRenderParam.renderHeight;
	//float pos[3] = {pixelIn[0], pixelIn[1], pixelIn[2]};	
}




extern "C"
void CurveRenderCUDA(RENDER_PARAM* multiSilRender, ScreenBlendVertex* dCurvePoint, unsigned char* surfaceIn, unsigned char* surfaceOut, size_t pitch, int PointRadius, bool IsForeshorten, bool BackCulling, bool BackDifferent)
{
	cudaMemcpyToSymbol(cRenderParam, multiSilRender, sizeof(RENDER_PARAM));

	CUDA_SAFE_CALL( cudaMemset(dSilOccupied,	  0, allPointSize) );
	CUDA_SAFE_CALL( cudaMemset(dSilOccupiedValid, 0, allPointSize) );
	dim3 Db = dim3(16,16);
	dim3 Dg = dim3( renderWidth/Db.x+1, renderHeight/Db.y+1 );

#if	UNIFORM_KNOT
	//pixelBits = 16;
	CUDA_SAFE_CALL( cudaMemset(dSilOccupied, 0, allPointSize) );
	CUDA_SAFE_CALL( cudaMemset(dSilOccupiedValid, 0, allPointSize) );
	GetCurveFlag<<<Dg,Db>>>((unsigned char*)surfaceIn, pitch, uint(renderWidth), uint(renderHeight), pixelBits, dSilOccupied, dSilOccupiedValid, asi);

	uint silNum;
	compact(&silNum, allPointSize, dNumber, dSilOccupied, dSilOccupiedValid, dOccupiedCompact, 0, NULL);  
	if (silNum>40000)
	{
		printf("Overflow of curve point!\n");
	}
	printf("%u initial guess!\n", silNum);
	printf("%u all point number!\n", allPointSize/sizeof(uint));


	//NewtonPolish<<<silNum/PIXELTHREADS+1, PIXELTHREADS>>>
	//	((unsigned char*)surfaceIn, (unsigned char*)surfaceOut, pitch,
	//	pixelBits,
	//	dOccupiedCompact, dSil, silNum, dFuncData, dPatchIndex,
	//	asi, degSize, 
	//	dDebugInfo
	//	);


	compactVertex<<<silNum/PIXELTHREADS+1, PIXELTHREADS>>>(pixelBits, dOccupiedCompact, dCurvePoint, silNum, IsForeshorten, PointRadius);

#else
#endif
}
